//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : reorder_kernels.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef REORDER_KERNELS_H
#define REORDER_KERNELS_H

namespace wavescomputer
{

extern __shared__ REAL sm[];

// forward
template <class T>
__global__ void kernel_reorder(T* dpsi, T* buffer)
{
    int tid = 3 * bx * Bx + tx;

    sm[tx]          = buffer[tid];
    sm[tx + Bx]     = buffer[tid + Bx];
    sm[tx + 2 * Bx] = buffer[tid + 2 * Bx];
    __syncthreads();

    dpsi[tx + bx * Bx]               = sm[3 * tx];
    dpsi[tx + bx * Bx + Gx * Bx]     = sm[3 * tx + 1];
    dpsi[tx + bx * Bx + 2 * Gx * Bx] = sm[3 * tx + 2];
}


// backward
template <class T>
__global__ void kernel_reorder2(T* buffer, T* dpsi)
{
    int tid = 3 * bx * Bx + tx;

    sm[3 * tx]     = dpsi[tx + bx * Bx];
    sm[3 * tx + 1] = dpsi[tx + bx * Bx + Gx * Bx];
    sm[3 * tx + 2] = dpsi[tx + bx * Bx + 2 * Gx * Bx];
    __syncthreads();

    buffer[tid]          = sm[tx];
    buffer[tid + Bx]     = sm[tx + Bx];
    buffer[tid + 2 * Bx] = sm[tx + 2 * Bx];
}

} // namespace wavescomputer

#endif // REORDER_KERNELS_H
